	.text
	.amdgcn_target "amdgcn-amd-amdhsa--gfx906"
	.protected	_Z23atomic_reduction_kernelPiS_i ; -- Begin function _Z23atomic_reduction_kernelPiS_i
	.globl	_Z23atomic_reduction_kernelPiS_i
	.p2align	8
	.type	_Z23atomic_reduction_kernelPiS_i,@function
_Z23atomic_reduction_kernelPiS_i:       ; @_Z23atomic_reduction_kernelPiS_i
; %bb.0:                                ; %entry
	s_add_u32 flat_scratch_lo, s12, s17
	s_addc_u32 flat_scratch_hi, s13, 0
	s_add_u32 s0, s0, s17
	s_load_dwordx4 s[48:51], s[8:9], 0x0
	s_load_dword s54, s[8:9], 0x10
	s_addc_u32 s1, s1, 0
	s_mov_b64 s[36:37], s[8:9]
	s_add_u32 s44, s36, 24
	v_lshlrev_b32_e32 v2, 20, v2
	v_lshlrev_b32_e32 v1, 10, v1
	s_mov_b32 s33, s16
	s_addc_u32 s45, s37, 0
	v_or3_b32 v40, v0, v1, v2
	s_mov_b32 s43, s14
	s_mov_b64 s[8:9], s[44:45]
	s_mov_b32 s12, s14
	s_mov_b32 s13, s15
	s_mov_b32 s14, s33
	v_mov_b32_e32 v31, v40
	v_mov_b32_e32 v0, 0
	s_mov_b32 s32, 0
	s_mov_b32 s42, s15
	s_mov_b64 s[34:35], s[10:11]
	s_mov_b64 s[38:39], s[6:7]
	s_mov_b64 s[40:41], s[4:5]
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, __ockl_get_group_id@rel32@lo+4
	s_addc_u32 s17, s17, __ockl_get_group_id@rel32@hi+12
	s_swappc_b64 s[30:31], s[16:17]
	v_mov_b32_e32 v42, v0
	s_mov_b64 s[4:5], s[40:41]
	s_mov_b64 s[6:7], s[38:39]
	s_mov_b64 s[8:9], s[44:45]
	s_mov_b64 s[10:11], s[34:35]
	s_mov_b32 s12, s43
	s_mov_b32 s13, s42
	s_mov_b32 s14, s33
	v_mov_b32_e32 v31, v40
	v_mov_b32_e32 v0, 0
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, __ockl_get_local_size@rel32@lo+4
	s_addc_u32 s17, s17, __ockl_get_local_size@rel32@hi+12
	s_swappc_b64 s[30:31], s[16:17]
	v_mov_b32_e32 v41, v0
	s_mov_b64 s[4:5], s[40:41]
	s_mov_b64 s[6:7], s[38:39]
	s_mov_b64 s[8:9], s[44:45]
	s_mov_b64 s[10:11], s[34:35]
	s_mov_b32 s12, s43
	s_mov_b32 s13, s42
	s_mov_b32 s14, s33
	v_mov_b32_e32 v31, v40
	v_mov_b32_e32 v0, 0
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, __ockl_get_local_id@rel32@lo+4
	s_addc_u32 s17, s17, __ockl_get_local_id@rel32@hi+12
	s_swappc_b64 s[30:31], s[16:17]
	v_mad_u64_u32 v[42:43], s[4:5], v41, v42, v[0:1]
	v_mov_b32_e32 v44, 0
	v_cmp_gt_i32_e32 vcc, s54, v42
	s_and_saveexec_b64 s[44:45], vcc
	s_cbranch_execz .LBB0_4
; %bb.1:                                ; %for.body.preheader
	s_mov_b64 s[46:47], 0
	v_mov_b32_e32 v44, 0
	s_getpc_b64 s[52:53]
	s_add_u32 s52, s52, __ockl_get_num_groups@rel32@lo+4
	s_addc_u32 s53, s53, __ockl_get_num_groups@rel32@hi+12
.LBB0_2:                                ; %for.body
                                        ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v43, 31, v42
	v_lshlrev_b64 v[0:1], 2, v[42:43]
	v_mov_b32_e32 v2, s49
	v_add_co_u32_e32 v0, vcc, s48, v0
	v_addc_co_u32_e32 v1, vcc, v2, v1, vcc
	global_load_dword v0, v[0:1], off
	s_add_u32 s8, s36, 24
	s_addc_u32 s9, s37, 0
	s_mov_b64 s[4:5], s[40:41]
	s_mov_b64 s[6:7], s[38:39]
	s_mov_b64 s[10:11], s[34:35]
	s_mov_b32 s12, s43
	s_mov_b32 s13, s42
	s_mov_b32 s14, s33
	v_mov_b32_e32 v31, v40
	s_waitcnt vmcnt(0)
	v_add_u32_e32 v44, v0, v44
	v_mov_b32_e32 v0, 0
	s_swappc_b64 s[30:31], s[52:53]
	v_mad_u64_u32 v[42:43], s[4:5], v0, v41, v[42:43]
	v_cmp_le_i32_e32 vcc, s54, v42
	s_or_b64 s[46:47], vcc, s[46:47]
	s_andn2_b64 exec, exec, s[46:47]
	s_cbranch_execnz .LBB0_2
; %bb.3:                                ; %Flow
	s_or_b64 exec, exec, s[46:47]
.LBB0_4:                                ; %Flow28
	s_or_b64 exec, exec, s[44:45]
	v_mov_b32_e32 v0, 0
	global_atomic_add v0, v44, s[50:51]
	s_endpgm
	.section	.rodata,#alloc
	.p2align	6, 0x0
	.amdhsa_kernel _Z23atomic_reduction_kernelPiS_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 16384
		.amdhsa_kernarg_size 80
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_private_segment_buffer 1
		.amdhsa_user_sgpr_dispatch_ptr 1
		.amdhsa_user_sgpr_queue_ptr 1
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 1
		.amdhsa_user_sgpr_flat_scratch_init 1
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_system_sgpr_private_segment_wavefront_offset 1
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 1
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 2
		.amdhsa_next_free_vgpr 45
		.amdhsa_next_free_sgpr 55
		.amdhsa_reserve_xnack_mask 1
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end0:
	.size	_Z23atomic_reduction_kernelPiS_i, .Lfunc_end0-_Z23atomic_reduction_kernelPiS_i
                                        ; -- End function
	.section	.AMDGPU.csdata
; Kernel info:
; codeLenInByte = 456
; NumSgprs: 61
; NumVgprs: 45
; ScratchSize: 16384
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 7
; VGPRBlocks: 11
; NumSGPRsForWavesPerEU: 61
; NumVGPRsForWavesPerEU: 45
; Occupancy: 5
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 1
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 1
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
	.text
	.protected	_Z24atomic_reduction_kernel2PiS_i ; -- Begin function _Z24atomic_reduction_kernel2PiS_i
	.globl	_Z24atomic_reduction_kernel2PiS_i
	.p2align	8
	.type	_Z24atomic_reduction_kernel2PiS_i,@function
_Z24atomic_reduction_kernel2PiS_i:      ; @_Z24atomic_reduction_kernel2PiS_i
; %bb.0:                                ; %entry
	s_add_u32 flat_scratch_lo, s12, s17
	s_addc_u32 flat_scratch_hi, s13, 0
	s_add_u32 s0, s0, s17
	s_load_dwordx4 s[48:51], s[8:9], 0x0
	s_load_dword s54, s[8:9], 0x10
	s_addc_u32 s1, s1, 0
	s_mov_b64 s[36:37], s[8:9]
	s_add_u32 s44, s36, 24
	v_lshlrev_b32_e32 v2, 20, v2
	v_lshlrev_b32_e32 v1, 10, v1
	s_mov_b32 s33, s16
	s_addc_u32 s45, s37, 0
	v_or3_b32 v40, v0, v1, v2
	s_mov_b32 s43, s14
	s_mov_b64 s[8:9], s[44:45]
	s_mov_b32 s12, s14
	s_mov_b32 s13, s15
	s_mov_b32 s14, s33
	v_mov_b32_e32 v31, v40
	v_mov_b32_e32 v0, 0
	s_mov_b32 s32, 0
	s_mov_b32 s42, s15
	s_mov_b64 s[34:35], s[10:11]
	s_mov_b64 s[38:39], s[6:7]
	s_mov_b64 s[40:41], s[4:5]
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, __ockl_get_group_id@rel32@lo+4
	s_addc_u32 s17, s17, __ockl_get_group_id@rel32@hi+12
	s_swappc_b64 s[30:31], s[16:17]
	v_mov_b32_e32 v42, v0
	s_mov_b64 s[4:5], s[40:41]
	s_mov_b64 s[6:7], s[38:39]
	s_mov_b64 s[8:9], s[44:45]
	s_mov_b64 s[10:11], s[34:35]
	s_mov_b32 s12, s43
	s_mov_b32 s13, s42
	s_mov_b32 s14, s33
	v_mov_b32_e32 v31, v40
	v_mov_b32_e32 v0, 0
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, __ockl_get_local_size@rel32@lo+4
	s_addc_u32 s17, s17, __ockl_get_local_size@rel32@hi+12
	s_swappc_b64 s[30:31], s[16:17]
	v_mov_b32_e32 v41, v0
	s_mov_b64 s[4:5], s[40:41]
	s_mov_b64 s[6:7], s[38:39]
	s_mov_b64 s[8:9], s[44:45]
	s_mov_b64 s[10:11], s[34:35]
	s_mov_b32 s12, s43
	s_mov_b32 s13, s42
	s_mov_b32 s14, s33
	v_mov_b32_e32 v31, v40
	v_mov_b32_e32 v0, 0
	v_mul_lo_u32 v42, v41, v42
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, __ockl_get_local_id@rel32@lo+4
	s_addc_u32 s17, s17, __ockl_get_local_id@rel32@hi+12
	s_swappc_b64 s[30:31], s[16:17]
	v_add_lshl_u32 v42, v42, v0, 4
	v_cmp_gt_i32_e32 vcc, s54, v42
	v_mov_b32_e32 v44, 0
	s_and_saveexec_b64 s[44:45], vcc
	s_cbranch_execz .LBB1_4
; %bb.1:                                ; %for.body.lr.ph
	v_lshlrev_b32_e32 v41, 4, v41
	s_mov_b64 s[46:47], 0
	v_mov_b32_e32 v44, 0
	s_getpc_b64 s[52:53]
	s_add_u32 s52, s52, __ockl_get_num_groups@rel32@lo+4
	s_addc_u32 s53, s53, __ockl_get_num_groups@rel32@hi+12
.LBB1_2:                                ; %for.body
                                        ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v43, 31, v42
	v_lshlrev_b64 v[0:1], 2, v[42:43]
	v_mov_b32_e32 v2, s49
	v_add_co_u32_e32 v16, vcc, s48, v0
	v_addc_co_u32_e32 v17, vcc, v2, v1, vcc
	global_load_dwordx4 v[0:3], v[16:17], off offset:16
	global_load_dwordx4 v[4:7], v[16:17], off
	global_load_dwordx4 v[8:11], v[16:17], off offset:48
	global_load_dwordx4 v[12:15], v[16:17], off offset:32
	s_add_u32 s8, s36, 24
	s_addc_u32 s9, s37, 0
	s_mov_b64 s[4:5], s[40:41]
	s_mov_b64 s[6:7], s[38:39]
	s_mov_b64 s[10:11], s[34:35]
	s_mov_b32 s12, s43
	s_mov_b32 s13, s42
	s_mov_b32 s14, s33
	v_mov_b32_e32 v31, v40
	s_waitcnt vmcnt(2)
	v_add_u32_e32 v3, v7, v3
	v_add_u32_e32 v2, v6, v2
	s_waitcnt vmcnt(0)
	v_add_u32_e32 v7, v15, v11
	v_add_u32_e32 v0, v4, v0
	v_add_u32_e32 v4, v14, v10
	v_add_u32_e32 v6, v12, v8
	v_add3_u32 v1, v5, v1, v3
	v_add3_u32 v3, v13, v9, v7
	v_add3_u32 v0, v0, v2, v1
	v_add3_u32 v1, v6, v4, v3
	v_add3_u32 v44, v0, v1, v44
	v_mov_b32_e32 v0, 0
	s_swappc_b64 s[30:31], s[52:53]
	v_mad_u64_u32 v[42:43], s[4:5], v41, v0, v[42:43]
	v_cmp_le_i32_e32 vcc, s54, v42
	s_or_b64 s[46:47], vcc, s[46:47]
	s_andn2_b64 exec, exec, s[46:47]
	s_cbranch_execnz .LBB1_2
; %bb.3:                                ; %Flow
	s_or_b64 exec, exec, s[46:47]
.LBB1_4:                                ; %Flow133
	s_or_b64 exec, exec, s[44:45]
	v_mov_b32_e32 v0, 0
	global_atomic_add v0, v44, s[50:51]
	s_endpgm
	.section	.rodata,#alloc
	.p2align	6, 0x0
	.amdhsa_kernel _Z24atomic_reduction_kernel2PiS_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 16384
		.amdhsa_kernarg_size 80
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_private_segment_buffer 1
		.amdhsa_user_sgpr_dispatch_ptr 1
		.amdhsa_user_sgpr_queue_ptr 1
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 1
		.amdhsa_user_sgpr_flat_scratch_init 1
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_system_sgpr_private_segment_wavefront_offset 1
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 1
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 2
		.amdhsa_next_free_vgpr 45
		.amdhsa_next_free_sgpr 55
		.amdhsa_reserve_xnack_mask 1
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end1:
	.size	_Z24atomic_reduction_kernel2PiS_i, .Lfunc_end1-_Z24atomic_reduction_kernel2PiS_i
                                        ; -- End function
	.section	.AMDGPU.csdata
; Kernel info:
; codeLenInByte = 552
; NumSgprs: 61
; NumVgprs: 45
; ScratchSize: 16384
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 7
; VGPRBlocks: 11
; NumSGPRsForWavesPerEU: 61
; NumVGPRsForWavesPerEU: 45
; Occupancy: 5
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 1
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 1
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
	.text
	.protected	_Z24atomic_reduction_kernel3PiS_i ; -- Begin function _Z24atomic_reduction_kernel3PiS_i
	.globl	_Z24atomic_reduction_kernel3PiS_i
	.p2align	8
	.type	_Z24atomic_reduction_kernel3PiS_i,@function
_Z24atomic_reduction_kernel3PiS_i:      ; @_Z24atomic_reduction_kernel3PiS_i
; %bb.0:                                ; %entry
	s_add_u32 flat_scratch_lo, s12, s17
	s_addc_u32 flat_scratch_hi, s13, 0
	s_add_u32 s0, s0, s17
	s_load_dwordx4 s[48:51], s[8:9], 0x0
	s_load_dword s54, s[8:9], 0x10
	s_addc_u32 s1, s1, 0
	s_mov_b64 s[36:37], s[8:9]
	s_add_u32 s44, s36, 24
	v_lshlrev_b32_e32 v2, 20, v2
	v_lshlrev_b32_e32 v1, 10, v1
	s_mov_b32 s33, s16
	s_addc_u32 s45, s37, 0
	v_or3_b32 v40, v0, v1, v2
	s_mov_b32 s43, s14
	s_mov_b64 s[8:9], s[44:45]
	s_mov_b32 s12, s14
	s_mov_b32 s13, s15
	s_mov_b32 s14, s33
	v_mov_b32_e32 v31, v40
	v_mov_b32_e32 v0, 0
	s_mov_b32 s32, 0
	s_mov_b32 s42, s15
	s_mov_b64 s[34:35], s[10:11]
	s_mov_b64 s[38:39], s[6:7]
	s_mov_b64 s[40:41], s[4:5]
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, __ockl_get_group_id@rel32@lo+4
	s_addc_u32 s17, s17, __ockl_get_group_id@rel32@hi+12
	s_swappc_b64 s[30:31], s[16:17]
	v_mov_b32_e32 v42, v0
	s_mov_b64 s[4:5], s[40:41]
	s_mov_b64 s[6:7], s[38:39]
	s_mov_b64 s[8:9], s[44:45]
	s_mov_b64 s[10:11], s[34:35]
	s_mov_b32 s12, s43
	s_mov_b32 s13, s42
	s_mov_b32 s14, s33
	v_mov_b32_e32 v31, v40
	v_mov_b32_e32 v0, 0
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, __ockl_get_local_size@rel32@lo+4
	s_addc_u32 s17, s17, __ockl_get_local_size@rel32@hi+12
	s_swappc_b64 s[30:31], s[16:17]
	v_mov_b32_e32 v41, v0
	s_mov_b64 s[4:5], s[40:41]
	s_mov_b64 s[6:7], s[38:39]
	s_mov_b64 s[8:9], s[44:45]
	s_mov_b64 s[10:11], s[34:35]
	s_mov_b32 s12, s43
	s_mov_b32 s13, s42
	s_mov_b32 s14, s33
	v_mov_b32_e32 v31, v40
	v_mov_b32_e32 v0, 0
	v_mul_lo_u32 v42, v41, v42
	s_getpc_b64 s[16:17]
	s_add_u32 s16, s16, __ockl_get_local_id@rel32@lo+4
	s_addc_u32 s17, s17, __ockl_get_local_id@rel32@hi+12
	s_swappc_b64 s[30:31], s[16:17]
	v_add_lshl_u32 v42, v42, v0, 2
	v_cmp_gt_i32_e32 vcc, s54, v42
	v_mov_b32_e32 v44, 0
	s_and_saveexec_b64 s[44:45], vcc
	s_cbranch_execz .LBB2_4
; %bb.1:                                ; %for.body.lr.ph
	v_lshlrev_b32_e32 v41, 2, v41
	s_mov_b64 s[46:47], 0
	v_mov_b32_e32 v44, 0
	s_getpc_b64 s[52:53]
	s_add_u32 s52, s52, __ockl_get_num_groups@rel32@lo+4
	s_addc_u32 s53, s53, __ockl_get_num_groups@rel32@hi+12
.LBB2_2:                                ; %for.body
                                        ; =>This Inner Loop Header: Depth=1
	v_ashrrev_i32_e32 v43, 31, v42
	v_lshlrev_b64 v[0:1], 2, v[42:43]
	v_mov_b32_e32 v2, s49
	v_add_co_u32_e32 v0, vcc, s48, v0
	v_addc_co_u32_e32 v1, vcc, v2, v1, vcc
	global_load_dwordx4 v[0:3], v[0:1], off
	s_add_u32 s8, s36, 24
	s_addc_u32 s9, s37, 0
	s_mov_b64 s[4:5], s[40:41]
	s_mov_b64 s[6:7], s[38:39]
	s_mov_b64 s[10:11], s[34:35]
	s_mov_b32 s12, s43
	s_mov_b32 s13, s42
	s_mov_b32 s14, s33
	v_mov_b32_e32 v31, v40
	s_waitcnt vmcnt(0)
	v_add_u32_e32 v1, v1, v3
	v_add_u32_e32 v0, v0, v2
	v_add3_u32 v44, v0, v1, v44
	v_mov_b32_e32 v0, 0
	s_swappc_b64 s[30:31], s[52:53]
	v_mad_u64_u32 v[42:43], s[4:5], v41, v0, v[42:43]
	v_cmp_le_i32_e32 vcc, s54, v42
	s_or_b64 s[46:47], vcc, s[46:47]
	s_andn2_b64 exec, exec, s[46:47]
	s_cbranch_execnz .LBB2_2
; %bb.3:                                ; %Flow
	s_or_b64 exec, exec, s[46:47]
.LBB2_4:                                ; %Flow50
	s_or_b64 exec, exec, s[44:45]
	v_mov_b32_e32 v0, 0
	global_atomic_add v0, v44, s[50:51]
	s_endpgm
	.section	.rodata,#alloc
	.p2align	6, 0x0
	.amdhsa_kernel _Z24atomic_reduction_kernel3PiS_i
		.amdhsa_group_segment_fixed_size 0
		.amdhsa_private_segment_fixed_size 16384
		.amdhsa_kernarg_size 80
		.amdhsa_user_sgpr_count 14
		.amdhsa_user_sgpr_private_segment_buffer 1
		.amdhsa_user_sgpr_dispatch_ptr 1
		.amdhsa_user_sgpr_queue_ptr 1
		.amdhsa_user_sgpr_kernarg_segment_ptr 1
		.amdhsa_user_sgpr_dispatch_id 1
		.amdhsa_user_sgpr_flat_scratch_init 1
		.amdhsa_user_sgpr_private_segment_size 0
		.amdhsa_system_sgpr_private_segment_wavefront_offset 1
		.amdhsa_system_sgpr_workgroup_id_x 1
		.amdhsa_system_sgpr_workgroup_id_y 1
		.amdhsa_system_sgpr_workgroup_id_z 1
		.amdhsa_system_sgpr_workgroup_info 0
		.amdhsa_system_vgpr_workitem_id 2
		.amdhsa_next_free_vgpr 45
		.amdhsa_next_free_sgpr 55
		.amdhsa_reserve_xnack_mask 1
		.amdhsa_float_round_mode_32 0
		.amdhsa_float_round_mode_16_64 0
		.amdhsa_float_denorm_mode_32 3
		.amdhsa_float_denorm_mode_16_64 3
		.amdhsa_dx10_clamp 1
		.amdhsa_ieee_mode 1
		.amdhsa_fp16_overflow 0
		.amdhsa_exception_fp_ieee_invalid_op 0
		.amdhsa_exception_fp_denorm_src 0
		.amdhsa_exception_fp_ieee_div_zero 0
		.amdhsa_exception_fp_ieee_overflow 0
		.amdhsa_exception_fp_ieee_underflow 0
		.amdhsa_exception_fp_ieee_inexact 0
		.amdhsa_exception_int_div_zero 0
	.end_amdhsa_kernel
	.text
.Lfunc_end2:
	.size	_Z24atomic_reduction_kernel3PiS_i, .Lfunc_end2-_Z24atomic_reduction_kernel3PiS_i
                                        ; -- End function
	.section	.AMDGPU.csdata
; Kernel info:
; codeLenInByte = 476
; NumSgprs: 61
; NumVgprs: 45
; ScratchSize: 16384
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 7
; VGPRBlocks: 11
; NumSGPRsForWavesPerEU: 61
; NumVGPRsForWavesPerEU: 45
; Occupancy: 5
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 1
; COMPUTE_PGM_RSRC2:USER_SGPR: 14
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 1
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 2
	.hidden	__ockl_get_group_id
	.hidden	__ockl_get_local_size
	.hidden	__ockl_get_local_id
	.hidden	__ockl_get_num_groups
	.ident	"clang version 16.0.6 (https://github.com/llvm/llvm-project.git 7cbf1a2591520c2491aa35339f227775f4d3adf6)"
	.ident	"AMD clang version 18.0.0git (https://github.com/RadeonOpenCompute/llvm-project roc-6.2.4 24392 1e2c94795ee0d6ab8e2ff3035965a6b74e11b475)"
	.section	".note.GNU-stack"
	.addrsig
	.amdgpu_metadata
---
amdhsa.kernels:
  - .args:
      - .address_space:  global
        .name:           in.coerce
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .name:           out.coerce
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .name:           ARRAYSIZE
        .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         32
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         40
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         48
        .size:           8
        .value_kind:     hidden_hostcall_buffer
      - .offset:         56
        .size:           8
        .value_kind:     hidden_default_queue
      - .offset:         64
        .size:           8
        .value_kind:     hidden_none
      - .offset:         72
        .size:           8
        .value_kind:     hidden_multigrid_sync_arg
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 80
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z23atomic_reduction_kernelPiS_i
    .private_segment_fixed_size: 16384
    .sgpr_count:     61
    .sgpr_spill_count: 0
    .symbol:         _Z23atomic_reduction_kernelPiS_i.kd
    .vgpr_count:     45
    .vgpr_spill_count: 0
    .wavefront_size: 64
  - .args:
      - .address_space:  global
        .name:           in.coerce
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .name:           out.coerce
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .name:           ARRAYSIZE
        .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         32
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         40
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         48
        .size:           8
        .value_kind:     hidden_hostcall_buffer
      - .offset:         56
        .size:           8
        .value_kind:     hidden_default_queue
      - .offset:         64
        .size:           8
        .value_kind:     hidden_none
      - .offset:         72
        .size:           8
        .value_kind:     hidden_multigrid_sync_arg
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 80
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z24atomic_reduction_kernel2PiS_i
    .private_segment_fixed_size: 16384
    .sgpr_count:     61
    .sgpr_spill_count: 0
    .symbol:         _Z24atomic_reduction_kernel2PiS_i.kd
    .vgpr_count:     45
    .vgpr_spill_count: 0
    .wavefront_size: 64
  - .args:
      - .address_space:  global
        .name:           in.coerce
        .offset:         0
        .size:           8
        .value_kind:     global_buffer
      - .address_space:  global
        .name:           out.coerce
        .offset:         8
        .size:           8
        .value_kind:     global_buffer
      - .name:           ARRAYSIZE
        .offset:         16
        .size:           4
        .value_kind:     by_value
      - .offset:         24
        .size:           8
        .value_kind:     hidden_global_offset_x
      - .offset:         32
        .size:           8
        .value_kind:     hidden_global_offset_y
      - .offset:         40
        .size:           8
        .value_kind:     hidden_global_offset_z
      - .offset:         48
        .size:           8
        .value_kind:     hidden_hostcall_buffer
      - .offset:         56
        .size:           8
        .value_kind:     hidden_default_queue
      - .offset:         64
        .size:           8
        .value_kind:     hidden_none
      - .offset:         72
        .size:           8
        .value_kind:     hidden_multigrid_sync_arg
    .group_segment_fixed_size: 0
    .kernarg_segment_align: 8
    .kernarg_segment_size: 80
    .language:       OpenCL C
    .language_version:
      - 2
      - 0
    .max_flat_workgroup_size: 1024
    .name:           _Z24atomic_reduction_kernel3PiS_i
    .private_segment_fixed_size: 16384
    .sgpr_count:     61
    .sgpr_spill_count: 0
    .symbol:         _Z24atomic_reduction_kernel3PiS_i.kd
    .vgpr_count:     45
    .vgpr_spill_count: 0
    .wavefront_size: 64
amdhsa.target:   amdgcn-amd-amdhsa--gfx906
amdhsa.version:
  - 1
  - 1
...

	.end_amdgpu_metadata
